Repository navigation
Fix more kernel launches over INT_MAX work-items with DPC++ 2026.2 - #3104
antonwolfy wants to merge 2 commits into
Conversation
With DPC++ 2026.2 the work-item ID range is recorded at the device link step, so modules built without the option at link time fail to launch kernels over more than INT_MAX work-items. Pass it at compile and link time to the ufunc, window and statistics extensions and to the dpnp_backend_c library, whose kernels can be launched over such ranges.
|
View rendered docs @ https://intelpython.github.io/dpnp/pull/3104/index.html |
|
Array API standard conformance tests for dpnp=0.21.0dev14=np2py314h8d9cdd5_6 ran successfully. |
|
No base build to compare — fix-id-queries-link-flag into master |
| target_compile_options(${python_module_name} PUBLIC -fno-sycl-id-queries-fit-in-int) | ||
| target_link_options( | ||
| ${python_module_name} | ||
| PUBLIC -fsycl-device-code-split=per_kernel -fno-sycl-id-queries-fit-in-int |
There was a problem hiding this comment.
blas extension also needs this flag
There was a problem hiding this comment.
For what? It has no pure SYCL kernel.
There was a problem hiding this comment.
It does have one.
syrk.cpp has its own copy_kernel that fills the lower triangle after oneMKL syrk: on GPU it is cgh.parallel_for(sycl::range<2>{n, n}, ...) at syrk.cpp:116, and syrk_impl always calls it (syrk.cpp:228).
So a.T @ a with a = dpnp.ones((1, 46341), dtype="f4") launches it over 46341 * 46341 > INT_MAX work-items, the same case this PR fixes for ufunc/window/statistics.
| * Fixed incorrect results of `dpnp.tensor.vecdot` in some cases with strided outputs and of `dpnp.tensor` reductions, `dpnp.tensor.vecdot` and `dpnp.tensor.matmul` on large inputs with some data types [#3082](https://github.com/IntelPython/dpnp/pull/3082) | ||
| * Fixed `dpnp.median` and `dpnp.nanmedian` raising a `ValueError` for a tuple `axis` when a kept dimension has size 0 [#3081](https://github.com/IntelPython/dpnp/pull/3081) | ||
| * Fixed a `RuntimeError` raised by `dpnp.tensor` kernels launched over more than `INT_MAX` work-items when built with DPC++ compiler 2026.2 or newer, by also passing `-fno-sycl-id-queries-fit-in-int` to the linker [#3089](https://github.com/IntelPython/dpnp/pull/3089) | ||
| * Fixed a `RuntimeError` raised by kernels of the `ufunc`, `window` and `statistics` extensions and of the legacy backend launched over more than `INT_MAX` work-items when built with DPC++ compiler 2026.2 or newer [#3104](https://github.com/IntelPython/dpnp/pull/3104) |
There was a problem hiding this comment.
there is a double space before the PR link
This branch has not been deployed
3 participants
This PR passes
-fno-sycl-id-queries-fit-in-intat both compile and link time to theufunc,windowandstatisticsextensions and to thedpnp_backend_clibrary. Until now only thedpnp.tensormodules and theindexingextension had it.With DPC++ 2026.2 the range a kernel assumes for its work-item IDs is recorded at the device link step, so these modules recorded the default
intrange, and their kernels failed when launched over more thanINT_MAXwork-items:The other extensions don't need the option:
fft,lapack,sparseandvmcompile no kernels of their own, and the onlyblaskernel (syrk) runs over ann x nrange, where no single dimension can exceedINT_MAX. Instatisticseach work-item processes several elements, so the limit is reached only for very large inputs.Validated with DPC++ 2026.2 on Intel Data Center GPU Max 1100 (Level Zero) and on CPU (OpenCL), on inputs with
2**31 + 1elements; results are the same on both devices:dpnp.gcdon broadcastint8inputsufuncdpnp.hamming(2**31 + 1)windowdpnp.random.logistic(size=2**31 + 1)dpnp_backend_ctest_binary_ufuncs.py,test_histogram.py,test_window.py,test_product.py,test_random_state.py,test_statistics.py,test_fft.pyandtest_special.pypass on both devices, and the build has no new warnings.No new tests are added: these cases need multi-GiB allocations, so as slow tests they would not run in CI and would skip on most devices.